#define WIDTH 1024          /* 宽 */
#define HEIGHT 1024         /* 高 */
#define WIDTH_MASK 1023     /* 宽的MASK */
#define HEIGHT_MASK 1023    /* 高的MASK */
#define WIDTH_POWER 10      /* 宽对2取对数  */
#define HEIGHT_POWER 10     /* 高对2取对数  */
#define CELLS_COUNT 1048576 /* 一共多少个点 */
#define V_FACTOR 0.96f      /* 速度的衰减因子 */

/*
    更新地图的内核
 */
#define MASK_SIX 63
#define getIdx(x, y) ((x & WIDTH_MASK) + ((y & HEIGHT_MASK) << WIDTH_POWER))
#define CELL_XY_MASK ~(((1<<12)-1)<<6)  /* 2+12个0  12个1  6个0 */
#define ENABLE_PRINT 1

constant float2 V_MIN = (float2) (-30, -30);
constant float2 V_MAX = (float2) ( 30,  30);

__kernel void vector_add(global const float *a, global const float *b,
                         global float *result) {
  int gid = get_global_id(0);
  result[gid] = a[gid] + b[gid];
}

/* 计算粒子的速度 */
__kernel void calc(global const float *ACCS, global int *MAP,
                      global const int *FMAP) {
  const int globalId = get_global_id(0);                  /* 全局第几个 */
  const int groupId = get_group_id(0);                    /* 第几组 */
  const int localSize = get_local_size(0);                /* 每组多少个 */
  const int localId = get_local_id(0);                    /* 组内第几个 */
  const int CELL_COL_WIDTH = WIDTH / get_num_groups(0);   /* 每组多宽 */
  const int CELL_X_START = CELL_COL_WIDTH * groupId;      /* 这组 WIDTH 的起点*/
  const int CELL_X_END = CELL_COL_WIDTH * (groupId + 1);  /* 这组 WIDTH 的终点*/

  if(ENABLE_PRINT && globalId==0) printf("OpenCL print, CALC, groupId=%d, localId=%d, globalId=%d, HEIGHT=%d, CELL_X_START=%d, CELL_X_END=%d", 
                            groupId, localId, globalId, HEIGHT, CELL_X_START, CELL_X_END);
  

  /* 逐行逐列计算 */
  for (int row = 0; row < HEIGHT; row++) {
    for (int col = CELL_X_START + localId; col < CELL_X_END;
         col += CELL_COL_WIDTH) {
      /* 小格子 */
      int cellIdx = getIdx(col, row);
      int cell    =   MAP[cellIdx];
      int type    =   (cell >> 30) & 3;
      int vx      =   (cell >> 24) & MASK_SIX - 32; 
      int vy      =   (cell >> 18) & MASK_SIX - 32; 
      int x       =   (cell >> 12) & MASK_SIX; 
      int y       =   (cell >>  6) & MASK_SIX; 
      int w       =   (cell >>  0) & MASK_SIX; 
      int safeW   =   max(w, 1);                 /* 除以 w 的时候，避免出现 0 */
      float2 a    =   (float2) (0.0, 0.0);
      float2 xy   =   (float2) (x, y);
      float2 vxy  =   (float2) (vx, vy);
      /* 如果不是空气，就计算附近 5x5 的粒子的影响 */
      if(type != 0) {
        for (int i = -2; i <= 2; i++) {
          for (int j = -2; j <= 2; j++) {
            if (i == 0 && j == 0) {
              continue;
            }
            int cell2 = MAP[getIdx(col + j, row + i)];
            if(cell2 != 0) {
              int type2   =   (cell2 >> 30) & 3;
              int x2      =   (cell2 >> 12) & MASK_SIX; 
              int y2      =   (cell2 >>  6) & MASK_SIX; 
              int w2      =   (cell2 >>  0) & MASK_SIX; 
              float2 xy2  =   (float2) (x2+i*64, y2+j*64);
              float2 dxy  = xy2 - xy;
              float dis   = length(dxy);
              a += ((w2 * ACCS[(int)dis]) / safeW) * normalize(dxy);
            }
          }
        }
        /* 更新速度 */
        vxy += a;
        vxy = clamp(vxy, V_MIN, V_MAX) * V_FACTOR;
        vx  = (int) vxy.x;
        vy  = (int) vxy.y;
        int ncell = (type<<30) | ((vx&MASK_SIX)<<24) | ((vy&MASK_SIX)<<18) | ((x&MASK_SIX)<<12) | ((y&MASK_SIX)<<6) | ((w&MASK_SIX)<<0);
        MAP[cellIdx] = ncell;
        if(ENABLE_PRINT && globalId==0 && cellIdx==0) printf("OpenCL print, CALC, cellIdx=%d, oldCell=%x, newCell=%x \n", cellIdx, cell, ncell);

      }
    }
  }
}

/* 移动粒子 */
__kernel void move(global int *MAP) {
  const int globalId = get_global_id(0);                  /* 全局第几个 */
  const int groupId = get_group_id(0);                    /* 第几组 */
  const int localSize = get_local_size(0);                /* 每组多少个 */
  const int localId = get_local_id(0);                    /* 组内第几个 */
  const int CELL_COL_WIDTH = WIDTH / get_num_groups(0);   /* 每组多宽 */
  const int CELL_X_START = CELL_COL_WIDTH * groupId;      /* 这组 WIDTH 的起点*/
  const int CELL_X_END = CELL_COL_WIDTH * (groupId + 1);  /* 这组 WIDTH 的终点*/

  if(ENABLE_PRINT && globalId==0) printf("OpenCL print, MOVE, groupId=%d, localId=%d, globalId=%d, HEIGHT=%d, CELL_X_START=%d, CELL_X_END=%d", 
                            groupId, localId, globalId, HEIGHT, CELL_X_START, CELL_X_END);

  /* 逐行逐列计算 */
  for (int row = 0; row < HEIGHT; row++) {
    for (int col = CELL_X_START + localId; col < CELL_X_END;
         col += CELL_COL_WIDTH) {
      /* 小格子 */
      int cellIdx = getIdx(col, row);
      int cell    =   MAP[cellIdx];
      int type    =   (cell >> 30) & 3;
      int vx      =   (cell >> 24) & MASK_SIX - 32; 
      int vy      =   (cell >> 18) & MASK_SIX - 32; 
      int x       =   (cell >> 12) & MASK_SIX; 
      int y       =   (cell >>  6) & MASK_SIX;
      int ncell, ntype, nx, ny, nvx, nvy;
      if(type == 0) {
        /* 空气，移入附近 3x3 的粒子 */
        for (int i = -1; i <= 1 && ncell==0; i++) {
          for (int j = -1; j <= 1 && ncell==0; j++) {
            if (i == 0 && j == 0) {
              continue;
            }
            int cell2 = MAP[getIdx(col + j, row + i)];
            if(cell2 != 0) {
              int type2   =   (cell2 >> 30) & 3;
              int vx2     =   (cell2 >> 24) & MASK_SIX - 32; 
              int vy2     =   (cell2 >> 18) & MASK_SIX - 32; 
              int x2      =   (cell2 >> 12) & MASK_SIX; 
              int y2      =   (cell2 >>  6) & MASK_SIX; 
              int nx2     =   x2 + vx2 + (64 * i);
              int ny2     =   y2 + vy2 + (64 * j);
              /* 如果新的位置在 [0,64) 内，则移动过来  */
              if (((nx2 | ny2) & (~MASK_SIX)) == 0) {
                nx = nx2;
                ny = ny2;
                ncell = cell2;
                ntype = type2;
                nvx = vx2;
                nvy = vy2;
              }
            }
          }
        }
      } else {
        /* 粒子考虑移出 */
        int nx2      =   x + vx;
        int ny2      =   y + vy;
        /* 是否还在本方格 */
        if (((nx2 | ny2) & (~MASK_SIX)) == 0) {
          ncell = cell;
          ntype = type;
          nx    = nx2;
          ny    = ny2;
          nvx   = vx;
          nvy   = vy;
        }
      }
      
      /* 计算新的元素 */
      if(ntype == 2) {
        /* 器灵需要计算伤害 */
        int vsum = (int) length((float2)(nvx, nvy));
        int hurt = min(24 - vsum, 0);
        ncell &= ~(((1<<12)-1)<<18);
        ncell |= (hurt&MASK_SIX)<<24;
      } else {
        /* 更新普通粒子 */
        ncell &= ~(((1<<12)-1)<<6);
        ncell |= (nx<<12) | (ny<<6);
      }

      MAP[cellIdx] = ncell;
      if(ENABLE_PRINT && globalId==0 && cellIdx==0) printf("OpenCL print, MOVE, cellIdx=%d, oldCell=%x, newCell=%x \n", cellIdx, cell, ncell);

      //MAP[cellIdx] = (cellIdx+1) | (1<<30);
    }
  }
}